<?xml version="1.0" encoding="UTF-8"?>
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.1 plus MathML 2.0//EN" "http://www.w3.org/Math/DTD/mathml2/xhtml-math11-f.dtd">
<html xmlns="http://www.w3.org/1999/xhtml" lang="en-US">
<head>
<meta http-equiv="Content-Type" content="text/html; charset=utf-8"/>
<meta name="viewport" content="width=device-width, initial-scale=0.4"/>
<meta name="google" content="notranslate" />
<link rel="canonical" href="https://sleef.org/quadcuda.xhtml" />
<link rel="icon" href="favicon.png" />
<link rel="stylesheet" type="text/css" href="texlike.css"/>
<link rel="stylesheet" type="text/css" href="//fonts.googleapis.com/css?family=Ubuntu" />
<link rel="stylesheet" type="text/css" href="sleef.css"/>
<title>SLEEF - Quad-precision math library reference (CUDA)</title>
</head>
<body translate="no" class="notranslate">
<h1>SLEEF - Quad-precision math library reference (CUDA)</h1>

<h2>Table of contents</h2>

<ul class="none" style="font-family: arial, sansserif; padding-left: 0.5cm;">
  <li><a class="underlined" href="index.xhtml">Introduction</a></li>
  <li><a class="underlined" href="compile.xhtml">Compiling and installing the library</a></li>
  <li><a class="underlined" href="purec.xhtml">Math library reference</a></li>
  <li>&nbsp;</li>
  <li><a class="underlined" href="quad.xhtml"> Quad-precision math library reference</a>
  <ul class="disc">
    <li><a class="underlined" href="quad.xhtml">Data types and functions for all architectures</a></li>
    <li><a class="underlined" href="quadx86.xhtml">Data types and functions for x86 architecture</a></li>
    <li><a class="underlined" href="quadaarch64.xhtml">Data types and functions for AArch64 architecture</a></li>
    <li><a class="underlined" href="quadppc64.xhtml">Data types and functions for PPC64 architecture</a></li>
    <li><a class="underlined" href="quads390x.xhtml">Data types and functions for System/390 architecture</a></li>
    <li><a class="underlined" href="quadcuda.xhtml">Data types and functions for CUDA</a></li>
    <ul class="circle">
      <li><a href="#tutorial">Tutorial</a></li>
      <li><a href="#conversion">Conversion functions</a></li>
      <li><a href="#comparison">Comparison and selection functions</a></li>
      <li><a href="#mathfunctions">Math functions</a></li>
    </ul>
  </ul>
  </li>
  <li>&nbsp;</li>
  <li><a class="underlined" href="dft.xhtml">DFT library reference</a></li>
  <li><a class="underlined" href="misc.xhtml">Other tools included in the package</a></li>
  <li><a class="underlined" href="benchmark.xhtml">Benchmark results</a></li>
  <li><a class="underlined" href="additional.xhtml">Additional notes</a></li>
</ul>


<h2 id="tutorial">Tutorial</h2>

<p class="noindent">
  Below is a <a class="underlined" href="hellocudaquad.cu">test
  code</a> for the CUDA functions. CUDA devices cannot directly
  compute with the QP FP data type. Thus, you have to
  use <b class="type">Sleef_quadx1</b> data type to retain a QP FP
  value in CUDA device codes. This data type has the same structure as
  the QP FP data type, and you can directly access the number by
  casting the pointer to the QP FP data type supported by the
  compiler. Beware of the strict-aliasing rule in this case.
</p>

<pre class="code">
<code>#include &lt;iostream&gt;</code>
<code>#include &lt;quadmath.h&gt;</code>
<code></code>
<code>#include &quot;sleefquadinline_cuda.h&quot;</code>
<code></code>
<code>// Based on the tutorial code at https://developer.nvidia.com/blog/even-easier-introduction-cuda/</code>
<code></code>
<code>__global__ void pow_gpu(int n, Sleef_quadx1 *r, Sleef_quadx1 *x, Sleef_quadx1 *y) {</code>
<code>  int index = threadIdx.x, stride = blockDim.x;</code>
<code></code>
<code>  for (int i = index; i &lt; n; i += stride)</code>
<code>    r[i] = Sleef_powq1_u10cuda(x[i], y[i]);</code>
<code>}</code>
<code></code>
<code>int main(void) {</code>
<code>  int N = 1 &lt;&lt; 20;</code>
<code></code>
<code>  Sleef_quadx1 *rd, *xd, *yd;</code>
<code>  cudaMallocManaged(&amp;rd, N*sizeof(Sleef_quadx1));</code>
<code>  cudaMallocManaged(&amp;xd, N*sizeof(Sleef_quadx1));</code>
<code>  cudaMallocManaged(&amp;yd, N*sizeof(Sleef_quadx1));</code>
<code></code>
<code>  __float128 *r = (__float128 *)rd, *x = (__float128 *)xd, *y = (__float128 *)yd;</code>
<code></code>
<code>  for (int i = 0; i &lt; N; i++) {</code>
<code>    r[i] = 0.0;</code>
<code>    x[i] = 1.00001Q;</code>
<code>    y[i] = i;</code>
<code>  }</code>
<code>  pow_gpu&lt;&lt;&lt;1, 256&gt;&gt;&gt;(N, rd, xd, yd);</code>
<code></code>
<code>  cudaDeviceSynchronize();</code>
<code></code>
<code>  double maxError = 0.0;</code>
<code>  for (int i = 0; i &lt; N; i++)</code>
<code>    maxError = fmax(maxError, fabsq(r[i]-powq(x[i], y[i])));</code>
<code>  std::cout &lt;&lt; &quot;Max error: &quot; &lt;&lt; maxError &lt;&lt; std::endl;</code>
<code></code>
<code>  cudaFree(yd);</code>
<code>  cudaFree(xd);</code>
<code>  cudaFree(rd);</code>
<code></code>
<code>  return 0;</code>
<code>}</code>
</pre>
<p style="text-align:center;">
  <a class="underlined" href="hellocudaquad.cu">Source code for testing CUDA functions</a>
</p>

<br/>
<br/>
<br/>

<p>
  You may want to use both CPU and GPU functions in the same source
  code. This is possible, as shown in <a class="underlined"
  href="hellocudaquad2.cu">the following test code</a>. You cannot use
  the library version of the SLEEF functions in CUDA source
  codes. Please include the header files for inlineable functions
  along with the header file for CUDA functions. The I/O functions are
  defined in sleefquadinline_purec_scalar.h.  You cannot
  use <b class="func">SLEEF_QUAD_C</b> or <b class="func">sleef_q</b>
  in device functions.
</p>

<pre class="code">
<code>// nvcc -O3 hellocudaquad2.cu -I./include --fmad=false -Xcompiler -ffp-contract=off</code>
<code></code>
<code>#include &lt;iostream&gt;</code>
<code>#include &lt;stdio.h&gt;</code>
<code>#include &lt;stdint.h&gt;</code>
<code>#include &lt;stdarg.h&gt;</code>
<code>#include &lt;ctype.h&gt;</code>
<code>#include &lt;assert.h&gt;</code>
<code>#include &lt;emmintrin.h&gt;</code>
<code></code>
<code>#include &quot;sleefquadinline_sse2.h&quot;</code>
<code>#include &quot;sleefquadinline_purec_scalar.h&quot;</code>
<code>#include &quot;sleefquadinline_cuda.h&quot;</code>
<code>#include &quot;sleefinline_sse2.h&quot;</code>
<code></code>
<code>// Based on the tutorial code at https://developer.nvidia.com/blog/even-easier-introduction-cuda/</code>
<code></code>
<code>__global__ void pow_gpu(int n, Sleef_quadx1 *r, Sleef_quadx1 *x, Sleef_quadx1 *y) {</code>
<code>  int index = threadIdx.x, stride = blockDim.x;</code>
<code></code>
<code>  for (int i = index; i &lt; n; i += stride)</code>
<code>    r[i] = Sleef_powq1_u10cuda(x[i], y[i]);</code>
<code>}</code>
<code></code>
<code>int main(void) {</code>
<code>  int N = 1 &lt;&lt; 20;</code>
<code></code>
<code>  Sleef_quadx1 *rd, *xd, *yd;</code>
<code>  cudaMallocManaged(&amp;rd, N*sizeof(Sleef_quadx1));</code>
<code>  cudaMallocManaged(&amp;xd, N*sizeof(Sleef_quadx1));</code>
<code>  cudaMallocManaged(&amp;yd, N*sizeof(Sleef_quadx1));</code>
<code></code>
<code>  Sleef_quad *r = (Sleef_quad *)rd, *x = (Sleef_quad *)xd, *y = (Sleef_quad *)yd;</code>
<code></code>
<code>  //</code>
<code></code>
<code>  for (int i = 0; i &lt; N; i++) {</code>
<code>    r[i] = Sleef_cast_from_doubleq1_purec(0);</code>
<code>    x[i] = Sleef_cast_from_doubleq1_purec(1.00001);</code>
<code>    y[i] = Sleef_cast_from_doubleq1_purec(i);</code>
<code>  }</code>
<code></code>
<code>  pow_gpu&lt;&lt;&lt;1, 256&gt;&gt;&gt;(N, rd, xd, yd);</code>
<code></code>
<code>  cudaDeviceSynchronize();</code>
<code></code>
<code>  Sleef_quadx2 maxError = Sleef_splatq2_sse2(Sleef_strtoq(&quot;0.0&quot;, NULL));</code>
<code></code>
<code>  for (int i = 0; i &lt; N; i += 2) {</code>
<code>    Sleef_quadx2 r2 = Sleef_loadq2_sse2(&amp;r[i]);</code>
<code>    Sleef_quadx2 x2 = Sleef_loadq2_sse2(&amp;x[i]);</code>
<code>    Sleef_quadx2 y2 = Sleef_loadq2_sse2(&amp;y[i]);</code>
<code></code>
<code>    Sleef_quadx2 q = Sleef_fabsq2_sse2(Sleef_subq2_u05sse2(r2, Sleef_powq2_u10sse2(x2, y2)));</code>
<code>    maxError = Sleef_fmaxq2_sse2(maxError, q);</code>
<code>  }</code>
<code></code>
<code>  Sleef_printf(&quot;Max error: %Qg\n&quot;,</code>
<code>               Sleef_fmaxq1_purec(Sleef_getq2_sse2(maxError, 0), Sleef_getq2_sse2(maxError, 1)));</code>
<code></code>
<code>  //</code>
<code></code>
<code>  cudaFree(yd);</code>
<code>  cudaFree(xd);</code>
<code>  cudaFree(rd);</code>
<code></code>
<code>  return 0;</code>
<code>}</code>
</pre>
<p style="text-align:center;">
  <a class="underlined" href="hellocudaquad2.cu">Source code for testing CUDA functions with CPU functions</a>
</p>


<h2 id="conversion">Conversion functions</h2>

<p class="funcname">Convert QP number to double-precision number</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">double</b> <b class="func">Sleef_cast_to_doubleq1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These functions convert a QP FP value to a double-precision value.
</p>

<hr/>

<p class="funcname">Convert double-precision number to QP number</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_cast_from_doubleq1_cuda</b>( <b class="type">double</b> <i class="var">a</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These functions convert a double-precision value to a QP FP value.
</p>

<hr/>


<p class="funcname">Convert QP number to 64-bit signed integer</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">int64_t</b> <b class="func">Sleef_cast_to_int64q1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These functions convert a QP FP value to a 64-bit signed integer.
</p>

<hr/>


<p class="funcname">Convert 64-bit signed integer to QP number</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_cast_from_int64q1_cuda</b>( <b class="type">int64_t</b> <i class="var">a</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These functions convert a 64-bit signed integer to a QP FP value.
</p>

<hr/>


<p class="funcname">Convert QP number to 64-bit unsigned integer</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">uint64_t</b> <b class="func">Sleef_cast_to_uint64q1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These functions convert a QP FP value to a 64-bit signed integer.
</p>

<hr/>


<p class="funcname">Convert 64-bit unsigned integer to QP number</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_cast_from_uint64q1_cuda</b>( <b class="type">uint64_t</b> <i class="var">a</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These functions convert a 64-bit unsigned integer to a QP FP value.
</p>


<h2 id="comparison">Comparison and selection functions</h2>

<p class="funcname">QP comparison functions</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">int32_t</b> <b class="func">Sleef_icmpltq1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i>, <b class="type">Sleef_quadx1</b> <i class="var">b</i> );<br/>
<b>__device__</b> <b class="type">int32_t</b> <b class="func">Sleef_icmpleq1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i>, <b class="type">Sleef_quadx1</b> <i class="var">b</i> );<br/>
<b>__device__</b> <b class="type">int32_t</b> <b class="func">Sleef_icmpgtq1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i>, <b class="type">Sleef_quadx1</b> <i class="var">b</i> );<br/>
<b>__device__</b> <b class="type">int32_t</b> <b class="func">Sleef_icmpgeq1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i>, <b class="type">Sleef_quadx1</b> <i class="var">b</i> );<br/>
<b>__device__</b> <b class="type">int32_t</b> <b class="func">Sleef_icmpeqq1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i>, <b class="type">Sleef_quadx1</b> <i class="var">b</i> );<br/>
<b>__device__</b> <b class="type">int32_t</b> <b class="func">Sleef_icmpneq1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i>, <b class="type">Sleef_quadx1</b> <i class="var">b</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions of <a class="underlined"
  href="quad.xhtml#basicComparison">comparison functions</a>.
</p>

<hr/>

<p class="funcname">QP comparison functions of the second kind</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">int32_t</b> <b class="func">Sleef_icmpq1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i>, <b class="type">Sleef_quadx1</b> <i class="var">b</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_icmpq1_purec"><b class="func">Sleef_icmpq1_purec</b></a>.
</p>

<hr/>

<p class="funcname">Check orderedness</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">int32_t</b> <b class="func">Sleef_iunordq1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i>, <b class="type">Sleef_quadx1</b> <i class="var">b</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_iunordq1_purec"><b class="func">Sleef_iunordq1_purec</b></a>.
</p>

<hr/>

<p class="funcname">Select elements</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_iselectq1_cuda</b>( <b class="type">int32_t</b> <i class="var">c</i>, <b class="type">Sleef_quadx1</b> <i class="var">a</i>, <b class="type">Sleef_quadx1</b> <i class="var">b</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions that operate in the same way as the ternary operator.
</p>


<h2 id="mathfunctions">Math functions</h2>

<p class="funcname">QP functions for basic arithmetic operations</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_addq1_u05cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i>, <b class="type">Sleef_quadx1</b> <i class="var">b</i> );<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_subq1_u05cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i>, <b class="type">Sleef_quadx1</b> <i class="var">b</i> );<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_mulq1_u05cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i>, <b class="type">Sleef_quadx1</b> <i class="var">b</i> );<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_divq1_u05cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i>, <b class="type">Sleef_quadx1</b> <i class="var">b</i> );<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_negq1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions of <a class="underlined"
  href="quad.xhtml#basicArithmetic">the basic arithmetic operations</a>.
</p>

<hr/>

<p class="funcname">Square root functions</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_sqrtq1_u05cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_sqrtq1_u05purec"><b class="func">Sleef_sqrtq1_u05purec</b></a>.
</p>

<hr/>

<p class="funcname">Sine functions</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_sinq1_u10cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_sinq1_u10purec"><b class="func">Sleef_sinq1_u10purec</b></a>.
</p>

<hr/>

<p class="funcname">Cosine functions</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_cosq1_u10cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_cosq1_u10purec"><b class="func">Sleef_cosq1_u10purec</b></a>.
</p>

<hr/>

<p class="funcname">Tangent functions</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_tanq1_u10cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_tanq1_u10purec"><b class="func">Sleef_tanq1_u10purec</b></a>.
</p>

<hr/>

<p class="funcname">Arc sine functions</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_asinq1_u10cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_asinq1_u10purec"><b class="func">Sleef_asinq1_u10purec</b></a>.
</p>

<hr/>

<p class="funcname">Arc cosine functions</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_acosq1_u10cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_acosq1_u10purec"><b class="func">Sleef_acosq1_u10purec</b></a>.
</p>

<hr/>

<p class="funcname">Arc tangent functions</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_atanq1_u10cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_atanq1_u10purec"><b class="func">Sleef_atanq1_u10purec</b></a>.
</p>

<hr/>

<p class="funcname">Base-<i>e</i> exponential functions</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_expq1_u10cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_expq1_u10purec"><b class="func">Sleef_expq1_u10purec</b></a>.
</p>

<hr/>

<p class="funcname">Base-2 exponential functions</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_exp2q1_u10cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_exp2q1_u10purec"><b class="func">Sleef_exp2q1_u10purec</b></a>.
</p>

<hr/>

<p class="funcname">Base-10 exponentail</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_exp10q1_u10cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_exp10q1_u10purec"><b class="func">Sleef_exp10q1_u10purec</b></a>.
</p>

<hr/>

<p class="funcname">Base-<i>e</i> exponential functions minus 1</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_expm1q1_u10cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_expm1q1_u10purec"><b class="func">Sleef_expm1q1_u10purec</b></a>.
</p>

<hr/>

<p class="funcname">Natural logarithmic functions</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_logq1_u10cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_logq1_u10purec"><b class="func">Sleef_logq1_u10purec</b></a>.
</p>

<hr/>

<p class="funcname">Base-2 logarithmic functions</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_log2q1_u10cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_log2q1_u10purec"><b class="func">Sleef_log2q1_u10purec</b></a>.
</p>

<hr/>

<p class="funcname">Base-10 logarithmic functions</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_log10q1_u10cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_log10q1_u10purec"><b class="func">Sleef_log10q1_u10purec</b></a>.
</p>

<hr/>

<p class="funcname">Logarithm of one plus argument</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_log1pq1_u10cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_log1pq1_u10purec"><b class="func">Sleef_log1pq1_u10purec</b></a>.
</p>

<hr/>

<p class="funcname">Power functions</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_powq1_u10cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">x</i>, <b class="type">Sleef_quadx1</b> <i class="var">y</i> );<br/>
<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_powq1_u10purec"><b class="func">Sleef_powq1_u10purec</b></a>.
</p>

<hr/>

<p class="funcname">Hyperbolic sine functions</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_sinhq1_u10cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_sinhq1_u10purec"><b class="func">Sleef_sinhq1_u10purec</b></a>.
</p>

<hr/>

<p class="funcname">Hyperbolic cosine functions</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_coshq1_u10cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_coshq1_u10purec"><b class="func">Sleef_coshq1_u10purec</b></a>.
</p>

<hr/>

<p class="funcname">Hyperbolic tangent functions</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_tanhq1_u10cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_tanhq1_u10purec"><b class="func">Sleef_tanhq1_u10purec</b></a>.
</p>

<hr/>

<p class="funcname">Inverse hyperbolic sine functions</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_asinhq1_u10cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_asinhq1_u10purec"><b class="func">Sleef_asinhq1_u10purec</b></a>.
</p>

<hr/>

<p class="funcname">Inverse hyperbolic cosine functions</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_acoshq1_u10cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_acoshq1_u10purec"><b class="func">Sleef_acoshq1_u10purec</b></a>.
</p>

<hr/>

<p class="funcname">Inverse hyperbolic tangent functions</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_atanhq1_u10cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_atanhq1_u10purec"><b class="func">Sleef_atanhq1_u10purec</b></a>.
</p>

<hr/>

<p class="funcname">Round to integer towards zero</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_truncq1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_truncq1_purec"><b class="func">Sleef_truncq1_purec</b></a>.
</p>

<hr/>

<p class="funcname">Round to integer towards minus infinity</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_floorq1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_floorq1_purec"><b class="func">Sleef_floorq1_purec</b></a>.
</p>

<hr/>

<p class="funcname">Round to integer towards plus infinity</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_ceilq1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_ceilq1_purec"><b class="func">Sleef_ceilq1_purec</b></a>.
</p>

<hr/>

<p class="funcname">Round to integer away from zero</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_roundq1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_roundq1_purec"><b class="func">Sleef_roundq1_purec</b></a>.
</p>

<hr/>

<p class="funcname">Round to integer, ties round to even</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_rintq1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_rintq1_purec"><b class="func">Sleef_rintq1_purec</b></a>.
</p>

<hr/>

<p class="funcname">Absolute value</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_fabsq1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">a</i> );<br/>
<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_fabsq1_purec"><b class="func">Sleef_fabsq1_purec</b></a>.
</p>

<hr/>

<p class="funcname">Copy sign of a number</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_copysignq1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">x</i>, <b class="type">Sleef_quadx1</b> <i class="var">y</i> );<br/>
<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_copysignq1_purec"><b class="func">Sleef_copysignq1_purec</b></a>.
</p>

<hr/>

<p class="funcname">Maximum of two numbers</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_fmaxq1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">x</i>, <b class="type">Sleef_quadx1</b> <i class="var">y</i> );<br/>
<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_fmaxq1_purec"><b class="func">Sleef_fmaxq1_purec</b></a>.
</p>

<hr/>

<p class="funcname">Minimum of two numbers</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_fminq1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">x</i>, <b class="type">Sleef_quadx1</b> <i class="var">y</i> );<br/>
<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_fminq1_purec"><b class="func">Sleef_fminq1_purec</b></a>.
</p>

<hr/>

<p class="funcname">Positive difference</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_fdimq1_u05cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">x</i>, <b class="type">Sleef_quadx1</b> <i class="var">y</i> );<br/>
<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_fdimq1_u05purec"><b class="func">Sleef_fdimq1_u05purec</b></a>.
</p>

<hr/>

<p class="funcname">Floating point remainder</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_fmodq1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">x</i>, <b class="type">Sleef_quadx1</b> <i class="var">y</i> );<br/>
<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_fmodq1_purec"><b class="func">Sleef_fmodq1_purec</b></a>.
</p>

<hr/>

<p class="funcname">Floating point remainder</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_remainderq1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">x</i>, <b class="type">Sleef_quadx1</b> <i class="var">y</i> );<br/>
<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_remainderq1_purec"><b class="func">Sleef_remainderq1_purec</b></a>.
</p>

<hr/>

<p class="funcname">Split a number to fractional and integral components</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_frexpq1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">x</i>, <b class="type">int32_t *</b> <i class="var">ptr</i> );<br/>
<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_frexpq1_purec"><b class="func">Sleef_frexpq1_purec</b></a>.
</p>

<hr/>

<p class="funcname">Break a number into integral and fractional parts</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_modfq1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">x</i>, <b class="type">Sleef_quadx1 *</b> <i class="var">ptr</i> );<br/>
<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_modfq1_purec"><b class="func">Sleef_modfq1_purec</b></a>.
</p>

<hr/>

<p class="funcname">2D Euclidian distance</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_hypotq1_u05cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">x</i>, <b class="type">Sleef_quadx1</b> <i class="var">y</i> );<br/>
<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_hypotq1_u05purec"><b class="func">Sleef_hypotq1_u05purec</b></a>.
</p>

<hr/>

<p class="funcname">Fused multiply and accumulate</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_fmaq1_u05cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">x</i>, <b class="type">Sleef_quadx1</b> <i class="var">y</i>, <b class="type">Sleef_quadx1</b> <i class="var">z</i> );<br/>
<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_fmaq1_u05purec"><b class="func">Sleef_fmaq1_u05purec</b></a>.
</p>

<hr/>

<p class="funcname">Multiply by integral power of 2</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">Sleef_quadx1</b> <b class="func">Sleef_ldexpq1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">x</i>, <b class="type">int32_t</b> <i class="var">e</i> );<br/>
<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_ldexpq1_purec"><b class="func">Sleef_ldexpq1_purec</b></a>.
</p>

<hr/>

<p class="funcname">Integer exponent of an FP number</p>

<p class="header">Synopsis</p>

<p class="synopsis">
#include &lt;sleefquadinline_cuda.h&gt;<br/>
<br/>
<b>__device__</b> <b class="type">int32_t</b> <b class="func">Sleef_ilogbq1_cuda</b>( <b class="type">Sleef_quadx1</b> <i class="var">x</i> );<br/>
<br/>
</p>

<p class="header">Description</p>

<p class="noindent">
  These are the vectorized functions
  of <a href="quad.xhtml#Sleef_ilogbq1_purec"><b class="func">Sleef_ilogbq1_purec</b></a>.
</p>


<p class="footer">
  Copyright &copy; 2010-2025 SLEEF Project, Naoki Shibata and contributors.<br/>
  SLEEF is open-source software and is distributed under the Boost Software License, Version 1.0.
</p>

<script type="text/javascript">
var sc_project=13098265; 
var sc_invisible=1; 
var sc_security="518de45e"; 
</script>
<script type="text/javascript"
src="https://www.statcounter.com/counter/counter.js"
async="async"></script>
<noscript><div class="statcounter"><a title="Web Analytics"
href="https://statcounter.com/" target="_blank"><img
class="statcounter"
src="https://c.statcounter.com/13098265/0/518de45e/1/"
alt="Web Analytics"
referrerPolicy="no-referrer-when-downgrade"></img></a></div></noscript>


</body>
</html>
